\chapter{Kod kernela}
\begin{program}
\caption{Plik Kernel.cl, część 1}
\begin{lstlisting}
__kernel void templateKernel(__global  uint * output,
                             __global  float * plA,
                             __global  float * plB,
                             __global  float * plC,
                             __global  float * plD,
                             __global  float * pt1x,
                             __global  float * pt1y,
                             __global  float * pt1z,
                             __global  float * pt2x,
                             __global  float * pt2y,
                             __global  float * pt2z,
                             __global  float * abso,
                             __local  uint * cache,
                             const     unsigned int level)
{
    uint L = get_global_id(0);
    uint LL = L;
    float Sx = 0.2; // harcoded values of source and receiver positions
    float Sy = 0.4;
    float Sz = 0.3;
    float Rx = 0.6;
    float Ry = -1.7;
    float Rz = 0.4;
    uint N = 6;
    bool czy_zly = false;
    bool czy_wektor = false;
\end{lstlisting}
\end{program}
\begin{program}
\caption{Plik Kernel.cl, część 2}
\begin{lstlisting}
    for(int i = 0 ; i<level; i++)
    {
        cache[i] = L%N;
        L=L/N;
        float t = -(plA[cache[i]]*Sx + plB[cache[i]]*Sy + plC[cache[i]]*Sz + plD[cache[i]])/(plA[cache[i]]*plA[cache[i]] + plB[cache[i]]*plB[cache[i]] + plC[cache[i]]*plC[cache[i]]);
        Sx = 2*plA[cache[i]]*t+Sx;
        Sy = 2*plB[cache[i]]*t+Sy;
        Sz = 2*plC[cache[i]]*t+Sz;
        
    }
    
    float vNx = Rx-Sx;
    float vNy = Ry-Sy;
    float vNz = Rz-Sz;

    for( int i = level-1 ; i>=0 ; i--)
    {
        float t = -(plA[cache[i]]*Sx + plB[cache[i]]*Sy + plC[cache[i]]*Sz + plD[cache[i]])/(plA[cache[i]]*vNx + plB[cache[i]]*vNy + plC[cache[i]]*vNz);
        
        float Cx = vNx*t + Sx;
        float Cy = vNy*t + Sy;
        float Cz = vNz*t + Sz;

\end{lstlisting}
\end{program}
\begin{program}
\caption{Plik Kernel.cl, część 3}
\begin{lstlisting}

        if( vNx*plA[cache[i]] + vNy*plB[cache[i]] + vNz*plC[cache[i]] < 0)
        {
            czy_wektor = true;
            break;
        }
        float eps = 0.00001;
        if(Cx - pt1x[cache[i]] < -eps || Cx - pt2x[cache[i]] > eps || Cy - pt1y[cache[i]] < -eps || Cy - pt2y[cache[i]] > eps || Cz - pt1z[cache[i]] < -eps || Cz - pt2z[cache[i]] > eps)
        {
            czy_zly = true;
            break;
        }
        float t2 = -(plA[cache[i]]*vNx + plB[cache[i]]*vNy + plC[cache[i]]*vNz)/(plA[cache[i]]*plA[cache[i]] + plB[cache[i]]*plB[cache[i]] + plC[cache[i]]*plC[cache[i]]);

        vNx = 2*plA[cache[i]]*t2 + vNx;
        vNy = 2*plB[cache[i]]*t2 + vNy;
        vNz = 2*plC[cache[i]]*t2 + vNz;
        Sx = Cx;
        Sy = Cy;
        Sz = Cz;
    }
    if(czy_zly == false && czy_wektor == false)
    {
        output[LL] = LL;
    }
    else
    {
        output[LL] = 0;
    }
}
\end{lstlisting}
\end{program}